Skip to content

[HIP] add test for uses of std::array on device code #244

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom
Merged
merged 1 commit into from
Jun 16, 2025
Merged

[HIP] add test for uses of std::array on device code #244

merged 1 commit into from
Jun 16, 2025

Conversation

jmmartinez
Copy link
Contributor

First attempt at testing llvm/llvm-project#136133 .

Draft for the moment, I have to manually test this further.

@jmmartinez jmmartinez requested review from Artem-B and yxsamliu May 13, 2025 11:50
@jmmartinez jmmartinez self-assigned this May 13, 2025
@jmmartinez jmmartinez mentioned this pull request May 13, 2025
Merged
yxsamliu
yxsamliu reviewed May 14, 2025
View reviewed changes
External/CUDA/array.cu Outdated
__global__ void kernel() { do_test(); }

int main() {
kernel<<<32, 32>>>();
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

can std::array be parameter type of kernel or device functions?

can it be used as non-constexpr stack variable in kernel or device functions?

if so, can we have testcases for those?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

can std::array be parameter type of kernel or device functions?

It can be used as a parameter for device and kernel functions (which would trigger the copy-constructor). According to the c++ reference std::array is an aggregate in the the same way as a struct with an array inside; so the same rules should apply. For example, all the constructors are implicitly declared.

I'm not sure that adding the copy-constructor case here will improve the coverage of the test (we're testing the interaction with the standard C++ library, and there is no code to test since these behaviors are implicitly defined by clang). Passing an aggregate as an argument should be covered by the tests on clang's side.

I can add the test for std::array constructors here too if you think we really need them.

can it be used as non-constexpr stack variable in kernel or device functions?

This is already covered in the test. We check std::array both in a constexpr and non-constexpr context.

yxsamliu
yxsamliu reviewed May 14, 2025
View reviewed changes
@jmmartinez jmmartinez force-pushed the cuda_wrappers/array branch 2 times, most recently from e9f5343 to 2731dd1 Compare May 16, 2025 14:15
@emankov emankov self-requested a review May 21, 2025 16:00
@jmmartinez jmmartinez force-pushed the cuda_wrappers/array branch from 2731dd1 to 631eac0 Compare June 16, 2025 07:58
@jmmartinez jmmartinez marked this pull request as ready for review June 16, 2025 07:58
@jmmartinez
Copy link
Contributor Author

I've moved the array test from CUDA to HIP since I have less control over CUDA buildbots. The hipify tests are already merged.

@jmmartinez
Copy link
Contributor Author

Would you be ok to:

  1. land this change and see how the build-bots behave.
  2. then land [CUDA][HIP] Add a __device__ version of std::__glibcxx_assert_fail() llvm-project#136133

@jmmartinez jmmartinez requested review from kzhuravl and yxsamliu June 16, 2025 08:02
@jmmartinez jmmartinez changed the title [CUDA][HIP] add test for uses of std::array on device code [HIP] add test for uses of std::array on device code Jun 16, 2025
@jmmartinez jmmartinez merged commit ecb8582 into llvm:main Jun 16, 2025
1 check passed
@Artem-B
Copy link
Member

Artem-B commented Jun 16, 2025

The commit appears to break CUDA tests:
https://lab.llvm.org/buildbot/#/builders/69/builds/22366 (see the "Configuring CUDA test-suite" section logs)

CMake Error at cmake/modules/TestSuite.cmake:48 (add_executable):
  No SOURCES given to target: array-cuda-12.4-c++20-libstdc++-9
Call Stack (most recent call first):
  cmake/modules/TestSuite.cmake:79 (llvm_test_executable_no_test)
  cmake/modules/GPUTestVariant.cmake:48 (llvm_test_executable)
  cmake/modules/GPUTestVariant.cmake:71 (create_one_local_test_f)
  External/CUDA/CMakeLists.txt:77 (create_one_local_test)
  External/CUDA/CMakeLists.txt:170 (create_local_cuda_tests)
  External/CUDA/CMakeLists.txt:347 (create_cuda_test_variant)
  External/CUDA/CMakeLists.txt:376 (create_cuda_tests)

Artem-B
Artem-B reviewed Jun 16, 2025
View reviewed changes
External/CUDA/CMakeLists.txt
@@ -63,6 +63,7 @@ macro(create_local_cuda_tests VariantSuffix)
list(APPEND CUDA_LOCAL_TESTS assert)
list(APPEND CUDA_LOCAL_TESTS axpy)
list(APPEND CUDA_LOCAL_TESTS algorithm)
list(APPEND CUDA_LOCAL_TESTS array)
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

There's no matching array.cu file committed under CUDA which makes cmake configuration fail.

@Artem-B Artem-B mentioned this pull request Jun 16, 2025
Merged
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Reviewers

@Artem-B Artem-B Artem-B left review comments

@yxsamliu yxsamliu yxsamliu approved these changes

@emankov emankov Awaiting requested review from emankov

@kzhuravl kzhuravl Awaiting requested review from kzhuravl

Assignees

@jmmartinez jmmartinez

Labels
None yet
Projects
None yet
Milestone
No milestone
Development

Successfully merging this pull request may close these issues.
4 participants
@jmmartinez @Artem-B @yxsamliu @kzhuravl